
#include "cudaBindings.h"

//-------------------
//--- BodiesOnGPU ---
//-------------------

float *BodiesOnGPU::copyToGPU(float *data_host, size_t count)
{
	void *res;
	size_t size = count * sizeof(float);
	SAFE_CALL(cudaMalloc((void **)&res, size));
	SAFE_CALL(cudaMemcpy(res, data_host, size, cudaMemcpyHostToDevice));
	return (float *)res;
}

void BodiesOnGPU::release(float *data_dev)
{
	SAFE_CALL(cudaFree(data_dev));
}

BodiesOnGPU::BodiesOnGPU(InterleavedBodies *bodies_host)
{
	pos_x = copyToGPU(bodies_host->positions.getX(),
					  bodies_host->getSize());
	pos_y = copyToGPU(bodies_host->positions.getY(),
					  bodies_host->getSize());
	pos_z = copyToGPU(bodies_host->positions.getZ(),
					  bodies_host->getSize());
	masses =  copyToGPU(&bodies_host->masses[0],
						bodies_host->getSize());
	count = bodies_host->getSize();
}

BodiesOnGPU::~BodiesOnGPU()
{
	release(pos_x);
	release(pos_y);
	release(pos_z);
	release(masses);
}

//----------------------
//--- cudaInteractor ---
//----------------------

__global__ void interactRange(float *pos_x, float *pos_y, float *pos_z,
							  float *masses, size_t body_count,
							  float *a_range_x, float *a_range_y, float *a_range_z,
							  size_t range_size, size_t offset)
{
	int n = blockIdx.x * blockDim.x + threadIdx.x;
	if (n < range_size)
	{
		//Copying current body.
		float px = pos_x[n + offset];
		float py = pos_y[n + offset];
		float pz = pos_z[n + offset];
		float mass = masses[n + offset];
		float ax = a_range_x[n + offset];
		float ay = a_range_y[n + offset];
		float az = a_range_z[n + offset];
		//Interacting.
		for (int i = 0; i < body_count; i++)
			if (i != n + offset)
				InteractionLogic::interactWith(px, py, pz,
											   ax, ay, az, mass,
											   pos_x[i], pos_y[i], pos_z[i],
											   masses[i]);
		//Storing acceleration.
		a_range_x[n] = ax;
		a_range_y[n] = ay;
		a_range_z[n] = az;
	}
}

void cudaInteractor::interactBodies(InterleavedBodies::SubVector bodies, size_t lo, size_t hi)
{
	//Copying accelerations of requred sub-vector to GPU.
	size_t size = (hi - lo) * sizeof(float);
	float *accelerations_dev_x;
	float *accelerations_dev_y;
	float *accelerations_dev_z;
	SAFE_CALL(cudaMalloc((void **)&accelerations_dev_x, size));
	SAFE_CALL(cudaMalloc((void **)&accelerations_dev_y, size));
	SAFE_CALL(cudaMalloc((void **)&accelerations_dev_z, size));
	SAFE_CALL(cudaMemcpy(accelerations_dev_x,
						 bodies.getOwner()->accelerations.getX() + lo,
						 size,
						 cudaMemcpyHostToDevice));
	SAFE_CALL(cudaMemcpy(accelerations_dev_y,
						 bodies.getOwner()->accelerations.getY() + lo,
						 size,
						 cudaMemcpyHostToDevice));
	SAFE_CALL(cudaMemcpy(accelerations_dev_z,
						 bodies.getOwner()->accelerations.getZ() + lo,
						 size,
						 cudaMemcpyHostToDevice));
	//Launching kernel.
	dim3 threads(512, 1);
	dim3 grid((hi - lo) / 512 + 1, 1);
	interactRange<<<threads, grid>>>(bodies_on_gpu->pos_x, bodies_on_gpu->pos_y,
									 bodies_on_gpu->pos_z, bodies_on_gpu->masses,
									 bodies_on_gpu->count,
									 accelerations_dev_x, accelerations_dev_y, accelerations_dev_z,
									 hi - lo, lo);
	CHECK_ERROR("\"interactRange\" kernel failed");
	//Copying result.
	SAFE_CALL(cudaMemcpy(bodies.getOwner()->accelerations.getX() + lo,
						 accelerations_dev_x,
						 size,
						 cudaMemcpyDeviceToHost));
	SAFE_CALL(cudaMemcpy(bodies.getOwner()->accelerations.getY() + lo,
						 accelerations_dev_y,
						 size,
						 cudaMemcpyDeviceToHost));
	SAFE_CALL(cudaMemcpy(bodies.getOwner()->accelerations.getZ() + lo,
						 accelerations_dev_z,
						 size,
						 cudaMemcpyDeviceToHost));
	SAFE_CALL(cudaFree(accelerations_dev_x));
	SAFE_CALL(cudaFree(accelerations_dev_y));
	SAFE_CALL(cudaFree(accelerations_dev_z));
}
